#pragma once

enum SpWalk {
  pull_fix, pull_dyn, push_seq, push_para, l2_opt, l1_opt
};

enum SpLoad {
  none, local, shared, shfl
};

enum DnCompute {
  u_e_v, u_v, e_v, u_e
};

enum Reduction {
  atomic, shfl_red, shared_red
};

template <typename F, typename Idx, typename DType,
          SpWalk spwalk, SpLoad spload, DnCompute dn_compute,
          Reduction reduction, uint sp_warp, uint dn_warp, 
          uint spipv, uint dnipv, uint dn_unroll>
void san_check() {
  if ( spipv > 32 ) {
    assert(reduction != shared_red);
    assert(spload != shared);
  } else {

  }
}


#define CALL_TUNER(spwalk, spload, dncompute, reduction, spg, dng, spv, dnv, dnit, tune_iters, ...)\
do{\
  for (int i=0; i<(tune_iters)/10; i++)\
  {\
    _spmm_manscript<F, int, float, spwalk, spload, dncompute, reduction, EdgeType::w, spg, dng, spv, dnv, dnit>\
    (__VA_ARGS__);\
  }\
  H_ERR(cudaDeviceSynchronize());\
  H_ERR(cudaEventRecord(s));\
  for (int i=0; i<iters; i++)\
  {\
    _spmm_manscript<F, int, float, spwalk, spload, dncompute, reduction, EdgeType::w, spg, dng, spv, dnv, dnit>\
    (__VA_ARGS__);\
  }\
  H_ERR(cudaEventRecord(e));\
  H_ERR(cudaDeviceSynchronize());\
  H_ERR(cudaEventElapsedTime(&time, s, e));\
  printf("spwalk=" #spwalk ",spload=" #spload ",dn_comp=" #dncompute\
        ",red=" #reduction ",sp:" #spg "v" #spv ",dn:" #dng "v" #dnv\
         "u" #dnit " :time=%.6f, TP=%.6f\n",\
  time/(tune_iters), (double)2*g.nedge*(tune_iters)*vlen/time/1e6);\
} while(0)